examples: a SYCL kernel and a HIP kernel behind the seam example 09 established - #575
Merged
Conversation
…stablished
Two examples, and the point of both is that they are example 09 one file name
apart. The seam, the CPU fallback, the constrained glob and the rule mechanism
are identical; what differs is which compiler consumes the island.
11-sycl-kernel is the one that tests the claim. Its island is ordinary C++ --
no dialect to see, no launch syntax -- and what makes it a device translation
unit is that it goes to a compiler with a device back end. It also carries the
two properties a SYCL build has and the others do not: TWO EDGES, because a
SYCL object holds its device image and nothing registers it until
`-fsycl-link` emits a wrapper; and TWO C++ RUNTIMES, because `libsycl.so` is
compiled against libstdc++ while an mcpp artifact links libc++. Example 09's
island can promise not to touch the standard library at all; this one cannot,
since SYCL is a C++ library, so the discipline moves from "no standard library"
to "nothing crosses the seam" -- the island catches its own `sycl::exception`
and returns a code.
12-hip-kernel is the same kernel with every device call spelled `hip*`. On the
NVIDIA platform HIP is a header layer over the CUDA runtime, so the compiler is
the project's own clang, `xim:hip-nvidia` contains no binaries, and the
`[dependencies.compat]` entry is the same one hop example 09 needs. Its
`accel` is two chunks -- `hip, cuda12.9+{sm_89}` -- and the second is character
for character what example 09 writes, because a device has one spelling in this
ecosystem however many programming models reach it.
Both answer `12 24 36 48` on an RTX 4080 and again under `--no-accel`, with
zero `/usr` paths on any command line.
The verification script gains three corrections a dry run made, two of them the
same shape -- a criterion pointed at the wrong object:
* `ldd` prints the program INTERPRETER with `=>` like everything else, so a
filter looking only for that arrow counted `/lib64/ld-linux-x86-64.so.2` as
a host library and failed a payload that was in fact complete.
* A program that cannot START leaves its loader error on the same stderr the
adapter check reads, so one defect was reported three times and the third
report named the wrong cause. The adapter checks are skipped when the
program did not start.
* A payload may be installed under either namespace, so the lookups accept
both -- the same two-namespace form xim-pkgindex's own tests use.
It also gains the assertion that would have caught the dpcpp regression: the
programs carry DT_RPATH and the payload's libraries carry no search path at
all, because a RUNPATH on one of those switches off the inherited RPATH of
whatever loaded it.
…t cannot Synchronised with the fixture the same island came from. Three things were needed to make "the island returns a code" true, and only two of them are a catch: the catch inside the buffer scope, the queue's asynchronous handler -- without which SYCL's default handler calls std::terminate, which no catch can intercept -- and, for the third, a build-time warning from the rule, because a missing device image is thrown from inside the SYCL scheduler and reaches neither path.
…only tested by the configuration that fails Two hypotheses were wrong before the backtrace settled it, and both were worth fixing anyway because each terminates on some machine. The actual cause reaches neither the caller's frame nor the queue's asynchronous handler, so the rule states it at build time instead.
…d take Four properties a consumer can rely on across all four lanes, each enforced by something rather than asserted: the device is spelled once, the engine knows no vendor name, no device command line reaches the host, and a payload is reachable rather than merely installed. Three open items with the measurement that would close each.
Found by the host verification run, not by building on this machine. `nvidia_hip_runtime_api.h` includes <cuda_profiler_api.h> at its second line and CUDA ships that header in its own component, so a machine with a host CUDA installation supplies it from /usr/include and the build works while depending on something it never declared. The fixture in mcpp-plugins gained the entry when CI refused it on a runner; the example is the same shape and had the same gap. V4 with this: 0 assertions failed -- both examples answer 12 24 36 48 on the device and again under --no-accel.
…it built against Two findings from the sandbox run, and the second is the more useful. EXAMPLE 11 STILL CARRIED THE PRE-2026.09.07 SHAPE: `compat:cuda-runtime` beside `compat:sycl-runtime`, plus the driver sentinel. The whole point of the adapter's 2026.09.07 version is that a SYCL project does not declare CUDA -- the SYCL runtime's back end dlopens the driver, so the adapter carries that hop. The fixture in mcpp-plugins was updated when the adapter was; the example was copied from it before that and kept the old form. SECTION G WAS READING A FARM THE EXAMPLE HAD NOT USED. It picked one out of the store by `sort -V | tail -1`, and a store that has seen two adapter versions holds two farms. In the sandbox that reported `ok` for every soname of the 2026.09.06 farm while the example under test built against a different one -- the same shape the round-3 notes record as "the report was chosen by `ls | head -1`", and the reason those notes exist. It now reads the directory the example's own `resolution.json` names, and says so out loud when it cannot: a run with no example build prints `note: falling back to the newest farm in the store` rather than quietly measuring something else. Two attempts were needed -- the first looked for `target/*/resolution.json` where the file is two levels down, which the note caught rather than hid. A property this exposed and the recipe comment gets wrong is recorded in a follow-up: `compat.sycl-runtime` has ONE `install()` for both versions and never reads `pkginfo.version()`, so 2026.09.06 does NOT "install the farm it always did" -- it installs whatever the current recipe installs. Only the descriptor entry is frozen. V4 after both fixes: 0 assertions failed, with section G naming the 2026.09.07 farm and no fallback.
…7.10 7.3 now carries what each row was decided by rather than that it was done. 7.9 lists the eight defects with the method that found each, because the method is the reusable part -- and names the shape four of them share: a criterion pointing at the wrong object, which then answers ok. The mechanical lesson is that a check which selects its own object must print which one it selected; section G now does, and that line is what turned the last of these from a false pass into a visible one. 7.10 records the one red this round did not cause: ci-fresh-install fails in apt-get inside a debian:11 container on five consecutive commits, two of them before this round. Diagnosed (the 404'd URL returns 200 from here, so a CDN node was at fault) and deliberately not patched, because it cannot be verified from this machine.
Sunrisepeak
pushed a commit
that referenced
this pull request
Sep 6, 2026
Both examples shipped in #575 and both verification runs measured them; I closed the surrounding rows and missed these two because the files were renamed from -compute to -kernel. Caught by counting the table's own statuses rather than rereading it.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
docs(plan): 7.8 -- the ecosystem review, and what each open item would take
Four properties a consumer can rely on across all four lanes, each enforced by
something rather than asserted: the device is spelled once, the engine knows no
vendor name, no device command line reaches the host, and a payload is
reachable rather than merely installed. Three open items with the measurement
that would close each.
docs(plan): 7.7 -- a promise a comment makes about error handling is only tested by the configuration that fails
Two hypotheses were wrong before the backtrace settled it, and both were worth
fixing anyway because each terminates on some machine. The actual cause reaches
neither the caller's frame nor the queue's asynchronous handler, so the rule
states it at build time instead.
examples: 11's island reports what it can, and its README says what it cannot
Synchronised with the fixture the same island came from. Three things were
needed to make "the island returns a code" true, and only two of them are a
catch: the catch inside the buffer scope, the queue's asynchronous handler --
without which SYCL's default handler calls std::terminate, which no catch can
intercept -- and, for the third, a build-time warning from the rule, because a
missing device image is thrown from inside the SYCL scheduler and reaches
neither path.
examples: a SYCL kernel and a HIP kernel behind the seam example 09 established
Two examples, and the point of both is that they are example 09 one file name
apart. The seam, the CPU fallback, the constrained glob and the rule mechanism
are identical; what differs is which compiler consumes the island.
11-sycl-kernel is the one that tests the claim. Its island is ordinary C++ --
no dialect to see, no launch syntax -- and what makes it a device translation
unit is that it goes to a compiler with a device back end. It also carries the
two properties a SYCL build has and the others do not: TWO EDGES, because a
SYCL object holds its device image and nothing registers it until
-fsycl-linkemits a wrapper; and TWO C++ RUNTIMES, becauselibsycl.soiscompiled against libstdc++ while an mcpp artifact links libc++. Example 09's
island can promise not to touch the standard library at all; this one cannot,
since SYCL is a C++ library, so the discipline moves from "no standard library"
to "nothing crosses the seam" -- the island catches its own
sycl::exceptionand returns a code.
12-hip-kernel is the same kernel with every device call spelled
hip*. On theNVIDIA platform HIP is a header layer over the CUDA runtime, so the compiler is
the project's own clang,
xim:hip-nvidiacontains no binaries, and the[dependencies.compat]entry is the same one hop example 09 needs. Itsaccelis two chunks --hip, cuda12.9+{sm_89}-- and the second is characterfor character what example 09 writes, because a device has one spelling in this
ecosystem however many programming models reach it.